#ifndef SYMMUL_COALESCE_KERNEL_CU
#define SYMMUL_COALESCE_KERNEL_CU
#include "symMul.h"
#include <cuda_runtime.h>

__global__ void symMulCoalesceKernel(float* M, float* B)
{
    int blockRow = blockIdx.y;
    int blockCol = blockIdx.x;

    float* Bsub;
    Bsub = &B[N*BLOCK_SIZE*blockRow+BLOCK_SIZE*blockCol];

    float Bvalue = 0;

    int row =threadIdx.y;
    int col =threadIdx.x;

    for (unsigned int k=0; k< GRID_SIZE; k++){	//Loop over all the blocks

        float* Msub = &M[N*BLOCK_SIZE*blockRow+BLOCK_SIZE*k];

        float* Mtsub = &M[N*BLOCK_SIZE*blockCol+BLOCK_SIZE*k];

        __shared__ float Ms[BLOCK_SIZE][BLOCK_SIZE];
        __shared__ float Mts[BLOCK_SIZE][BLOCK_SIZE];

        Ms[row][col]= Msub[row*N+col];
        Mts[row][col]= Mtsub[row*N+col];

        __syncthreads();

        for (int e=0; e<BLOCK_SIZE; e++){
            Bvalue += Ms[row][e]*Mts[col][e];
        }

        __syncthreads();
    }

    Bsub[idx(row,col)] = Bvalue;
}

#endif
